Fix nonzero bool nonstandard bytes - #3055
Conversation
`NonZeroIndicator` compared against `inputT(0)` and the `Cumsum1D` factories used `NoOpTransformer`; for `bool` both fold into a raw byte load, so the scan summed byte values instead of 0/1. A mask stored as [0, 1, 2, 255, 0, 1] reported 259 non-zeros instead of 4, affecting `nonzero`, `where`, `extract`, `place` and `repeat`. Cast bool via `sycl::bit_cast<std::uint8_t>` / `CastTransformer`, as `convert_impl` already does (IntelPythongh-2121).
C++ may fold a bool comparison into a raw byte load, so a byte other than 0x00/0x01 compared unequal to a normalized True, ordered by its byte value, and leaked into computed bool results. Such a byte arises when a buffer is written through a raw pointer or viewed from integer data. Add `normalize_bool` and apply it where a bool is read from memory: elementwise operand loads (bool is excluded from the vector paths, which cannot normalize per element), the `convert_impl` same-type branch, the search-reduction loads, the `isin` equality test and the argsort projection. Add bool comparators for the merge-sort path. `sort` now orders False before True rather than reproducing NumPy's raw byte order, which would carry garbage bytes through a sort.
|
Can one of the admins verify this patch? |
ndgrigorian
left a comment
ndgrigorian
left a comment
There was a problem hiding this comment.
I've looked over the changes, seems like a welcome fix for a previously unnoticed bug, LGTM
|
@abagusetty seems that the tests fail with the open-source compiler, interestingly enough. Possible bug in the nightly DPC++? |
|
coverage: 78.472% (-0.02%) from 78.491% — abagusetty:fix-nonzero-bool-nonstandard-bytes into IntelPython:master |
`bool` takes a single radix pass, so the bucket index is `byte & 0xF`. An unnormalized byte whose low nibble is zero (0x10, 0x80, 0xF0) bucketed as False and sorted before True elements. Normalize in `order_preserving_cast`, where every radix path converges, and take the argument by reference so a copy cannot let the compiler assume a 0/1 byte.
No bug in nightly. Interestingly, it is doing great. The issue was UB in the PR that the two compilers are treating it differently. A bool whose byte isnt |
|
|
||
| using dpnp::tensor::sycl_utils::sub_group_load; | ||
| using dpnp::tensor::sycl_utils::sub_group_store; | ||
| using dpnp::tensor::type_utils::normalize_bool; |
There was a problem hiding this comment.
It seems there is one more gap with BitwiseXorInplaceFunctor in dpnp/tensor/libtensor/include/kernels/elementwise_functions/bitwise_xor.hpp:300:
if constexpr (std::is_same_v<resT, bool>) {
res = (res != in);
}The in-place functor normalizes only the RHS (op(lhs[k], normalize_bool(rhs[k]))); res/lhs is read raw; != is value-use, so a non-canonical destination byte gives the wrong answer:
import numpy as np, dpnp
def bviews(u8): # (numpy bool view, dpnp bool view) of same bytes
a = np.array(u8, dtype=np.uint8)
return a.view(bool), dpnp.asarray(a).view(dpnp.bool)
# non-canonical lhs (0x02=True), canonical rhs (0x01=True); True ^ True = False
na, da = bviews([2])
nb, db = bviews([1])
# out-of-place xor: PR normalizes BOTH inputs -> expect PASS on PR branch
print("a ^ b numpy:", np.bitwise_xor(na, nb),
" dpnp:", dpnp.asnumpy(dpnp.bitwise_xor(da, db)))
# Out: a ^ b numpy: [False] dpnp: [False]
# in-place xor: PR normalizes only rhs, lhs read raw via (res != in) -> expect FAIL
na2, da2 = bviews([2]); _, db2 = bviews([1])
na2 ^= nb
da2 ^= db2
print("a ^= b numpy:", na2, " dpnp:", dpnp.asnumpy(da2))
# Out: a ^= b numpy: [False] dpnp: [ True]So we probably need to update BitwiseXorInplaceFunctor with:
if constexpr (std::is_same_v<resT, bool>) {
res = (normalize_bool(res) != in); // in already normalized by caller; keeps output canonical
}and to cover with a test.
| // takes by reference: copying a bool first would let the compiler assume a | ||
| // 0/1 byte and fold the normalization away, see gh-2121 |
There was a problem hiding this comment.
The comment does not seem correct for value merge/insertion sort, which copies elements into auto locals before the comparator runs (merge_sort.hpp lines ~83–84, ~228).
The reference binds to those copies, not memory. It's harmless (the in-comparator normalize_bool is what makes sort correct, and CI's sort tests pass), but the rationale is wrong and only true for the argsort path.
We should either correct or drop the comment.
5 participants
Boolean arrays may contain non-zero bytes like
0x02or0xFFthat NumPy treats as True but dpnp does not.Fixes #3054